extern "C"
__global__ void fft(int size, int N, float* a){
	__shared__ float w_list[WLIST_SIZE*2];
	int step;
	int step2;
	int flag = 0;
 	int dsInd;
	int raInd;
	int rbInd;
	int rsaInd;
	int rsbInd;
	int wiInd;
	int wjInd;
	
	float w1[2];
	float w2[2];
	float tmp[2];
	float tmp2[2];
	
	float a00a;
	float a00b;
	float a01a;
	float a01b;
	float a10a;
	float a10b;
	float a11a;
	float a11b;
	
	int a00ai;
	int a00bi;
	int a01ai;
	int a01bi;
	int a10ai;
	int a10bi;
	int a11ai;
	int a11bi;
	
	for (int k = threadIdx.y * blockDim.x + threadIdx.x; k < N; k = k + blockDim.x * blockDim.y) {
		a00a = -2*3.14159265358979323846*k/N;
		w_list[2*k] = cos(a00a);
		w_list[2*k+1] = sin(a00a);
	}
	__syncthreads();
	
	for (int q = 0; q < size; q++) {
		step = ceil(pow(2.0f,q));
		step2 = step*2;
		for (int down = threadIdx.y; down < N - step; down = down + blockDim.y) {
			for (int right = threadIdx.x; right < N - step; right = right + blockDim.x) {
				flag = (down/step) % 2;
				flag = flag + (right/step) % 2;
				if (flag == 0){
					wiInd = (down % step2)*(N/step2) * 2;
					wjInd = (right % step2)*(N/step2) * 2;
					w1[0] = w_list[wiInd];
					w1[1] = w_list[wiInd+1];
					w2[0] = w_list[wjInd];
					w2[1] = w_list[wjInd+1];
					
					dsInd = down + step;
					raInd = 2*right;
					rbInd = raInd + 1;
					rsaInd = 2*(right+step);
					rsbInd = rsaInd+1;
					
					a00ai = down*N*2 + raInd;
					a00bi = a00ai + 1;
					a01ai = down*N*2 + rsaInd;
					a01bi = a01ai + 1;
					a10ai = dsInd*N*2 + raInd;
					a10bi = a10ai + 1;
					a11ai = dsInd*N*2 + rsaInd;
					a11bi = a11ai + 1;
										
					
					//XXX = XXX.add(a00);
					a00a = a[a00ai];
					a00b = a[a00bi];
					a10a = a[a00ai];
					a10b = a[a00bi];
					a01a = a[a00ai];
					a01b = a[a00bi];
					a11a = a[a00ai];
					a11b = a[a00bi];
					
					
					//XXX = XXX.add(t1.multy(a10));
					tmp[0] = w1[0]*a[a10ai] - w1[1]*a[a10bi];
					tmp[1] = w1[1]*a[a10ai] + w1[0]*a[a10bi];
					a00a = a00a + tmp[0];
					a00b = a00b + tmp[1];
					a10a = a10a - tmp[0];
					a10b = a10b - tmp[1];
					a01a = a01a + tmp[0];
					a01b = a01b + tmp[1];
					a11a = a11a - tmp[0];
					a11b = a11b - tmp[1];
					
					
					//XXX = XXX.add(t2.multy(a01));
					tmp[0] = w2[0]*a[a01ai] - w2[1]*a[a01bi];
					tmp[1] = w2[1]*a[a01ai] + w2[0]*a[a01bi];
					a00a = a00a + tmp[0];
					a00b = a00b + tmp[1];
					a10a = a10a + tmp[0];
					a10b = a10b + tmp[1];
					a01a = a01a - tmp[0];
					a01b = a01b - tmp[1];
					a11a = a11a - tmp[0];
					a11b = a11b - tmp[1];
					
					
					//XXX = XXX.add(t1.multy(t2.multy(a11)));
					tmp[0] = w2[0]*a[a11ai] - w2[1]*a[a11bi];
					tmp[1] = w2[1]*a[a11ai] + w2[0]*a[a11bi];
					tmp2[0] = w1[0]*tmp[0] - w1[1]*tmp[1];
					tmp2[1] = w1[0]*tmp[1] + w1[1]*tmp[0];
					
					a00a = a00a + tmp2[0];
					a00b = a00b + tmp2[1];
					a10a = a10a - tmp2[0];
					a10b = a10b - tmp2[1];
					a01a = a01a - tmp2[0];
					a01b = a01b - tmp2[1];
					a11a = a11a + tmp2[0];
					a11b = a11b + tmp2[1];
					
					a[a00ai] = a00a;
					a[a01ai] = a01a;
					a[a10ai] = a10a;
					a[a11ai] = a11a;
					a[a00bi] = a00b;
					a[a01bi] = a01b;
					a[a10bi] = a10b;
					a[a11bi] = a11b;
					
				}	
			}
		}
		__syncthreads();
	}
}